[Backend] Refactor Transform Pipeline to support different backends#2189
Conversation
|
Note Reviews pausedIt looks like this branch is under active development. To avoid overwhelming you with review comments due to an influx of new commits, CodeRabbit has automatically paused this review. You can configure this behavior by changing the Use the following commands to manage reviews:
Use the checkboxes below for quick actions:
No actionable comments were generated in the recent review. 🎉 ℹ️ Recent review info⚙️ Run configurationConfiguration used: defaults Review profile: CHILL Plan: Pro Run ID: 📒 Files selected for processing (6)
🚧 Files skipped from review as they are similar to previous changes (1)
📝 WalkthroughWalkthroughAdds a Pipeline registry and per-backend pipeline bodies (CPU/CUDA/ROCm/Metal), pass-context helpers, a PreLowerSemanticCheck pass, refactors engine and JIT lowering to resolve and invoke target-specific pipelines, and updates CUDA-targeted tests to use the new CUDA prologue lowering. ChangesBackend Pipeline Architecture and Integration
Estimated code review effort🎯 4 (Complex) | ⏱️ ~60 minutes Possibly related PRs
Suggested reviewers
Poem
🚥 Pre-merge checks | ✅ 4 | ❌ 1❌ Failed checks (1 warning)
✅ Passed checks (4 passed)
✏️ Tip: You can configure your own custom pre-merge checks in the settings. ✨ Finishing Touches🧪 Generate unit tests (beta)
Thanks for using CodeRabbit! It's free for OSS, and your support helps us grow. If you like it, consider giving us a shout-out. Comment |
|
👋 Hi! Thank you for contributing to the TileLang project. Please remember to run We appreciate you taking this step! Our team will review your contribution, and we look forward to your awesome work! 🚀 |
3241ead to
997c2e2
Compare
997c2e2 to
f83a7e1
Compare
Introduce Pipeline abstraction in backend/pipeline.py with per-backend registration. Each backend (cuda, hip, c, llvm) now registers its own compilation pass pipeline. engine/lower.py resolves the pipeline from the target instead of hardcoding phase imports.
05cde9d to
a267e6e
Compare
a267e6e to
523136f
Compare
There was a problem hiding this comment.
Actionable comments posted: 2
🧹 Nitpick comments (1)
tilelang/backend/rocm/pipeline.py (1)
18-82: 🏗️ Heavy liftConsider extracting a shared lowering spine with backend hook points.
This function duplicates most of
tilelang.engine.pass_pipeline.LowerCommon(ordering and pass set), which will be costly to keep in sync as pipelines evolve. A shared helper with backend-specific pre/post hooks would reduce drift risk.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tilelang/backend/rocm/pipeline.py` around lines 18 - 82, The lower_amd function duplicates the bulk of tilelang.engine.pass_pipeline.LowerCommon; refactor by extracting the shared lowering spine into a single helper (e.g., LowerCommonSpine or reuse LowerCommon) that encapsulates the common sequence of transforms and returns a mod and pass_ctx, then change lower_amd to call that helper and only perform AMD-specific pre/post hook passes (e.g., any passes before/after LayoutVisual, the calls that rely on pass_ctx like VectorizeLoop(enable_vectorize=allow_vectorize(pass_ctx=pass_ctx)) and MergeSharedMemoryAllocations(...should_enable_aggressive_merge(pass_ctx=pass_ctx)), and conditional branches like allow_global_thread_synchronization()); keep unique AMD-only transforms (e.g., LowerLDGSTG, ThreadSync("shared.dyn"), MarkCudaSyncCalls(False), LowerDeviceKernelLaunch) in lower_amd and delegate the rest to the shared helper so ordering stays centralized.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@tilelang/backend/pipeline.py`:
- Around line 29-36: The register_pipeline function currently overwrites
existing entries in the _PIPELINES mapping when pipeline.name already exists;
change register_pipeline to first check if pipeline.name is already present in
_PIPELINES and, instead of silently replacing it, raise a clear exception (e.g.,
ValueError) or log/error and refuse registration so callers must explicitly
unregister/replace before adding; reference the register_pipeline function, the
_PIPELINES mapping, and pipeline.name when locating where to add the existence
check and error handling.
In `@tilelang/engine/pass_pipeline.py`:
- Around line 68-85: The code treats "all" specially only when formats_str ==
"all", so mixed inputs like "txt,all" end up forwarding the literal "all";
update the parsing logic so after building formats_list from formats_str (the
variable formats_list), if "all" appears expand/replace it with the canonical
list ["txt","png","pdf","svg"] (or return that union) before performing the
invalid_formats check; ensure valid_formats used for validation does not treat
"all" as a valid concrete format (keep "all" only as a special token) and
validate against the canonical formats, referencing formats_str, formats_list,
valid_formats, invalid_formats and the TL_LAYOUT_VISUALIZATION_FORMATS setting.
---
Nitpick comments:
In `@tilelang/backend/rocm/pipeline.py`:
- Around line 18-82: The lower_amd function duplicates the bulk of
tilelang.engine.pass_pipeline.LowerCommon; refactor by extracting the shared
lowering spine into a single helper (e.g., LowerCommonSpine or reuse
LowerCommon) that encapsulates the common sequence of transforms and returns a
mod and pass_ctx, then change lower_amd to call that helper and only perform
AMD-specific pre/post hook passes (e.g., any passes before/after LayoutVisual,
the calls that rely on pass_ctx like
VectorizeLoop(enable_vectorize=allow_vectorize(pass_ctx=pass_ctx)) and
MergeSharedMemoryAllocations(...should_enable_aggressive_merge(pass_ctx=pass_ctx)),
and conditional branches like allow_global_thread_synchronization()); keep
unique AMD-only transforms (e.g., LowerLDGSTG, ThreadSync("shared.dyn"),
MarkCudaSyncCalls(False), LowerDeviceKernelLaunch) in lower_amd and delegate the
rest to the shared helper so ordering stays centralized.
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: 425c0493-ca4b-4cec-bf94-f136d7920f16
📒 Files selected for processing (16)
testing/python/transform/_transform_testing_utils.pytesting/python/transform/test_tilelang_transform_inject_tcgen05_fence.pytesting/python/transform/test_tilelang_transform_lexical_alloc_scope.pytesting/python/transform/test_tilelang_transform_lower_shared_barrier.pytesting/python/transform/test_tilelang_transform_plan_update_buffer_allocation_location.pytilelang/backend/__init__.pytilelang/backend/common.pytilelang/backend/cpu/__init__.pytilelang/backend/cuda/__init__.pytilelang/backend/cuda/pipeline.pytilelang/backend/pipeline.pytilelang/backend/rocm/__init__.pytilelang/backend/rocm/pipeline.pytilelang/engine/lower.pytilelang/engine/pass_pipeline.pytilelang/jit/adapter/utils.py
| def register_pipeline(pipeline: Pipeline) -> Pipeline: | ||
| """Register a lowering pipeline for a backend. | ||
|
|
||
| The pipeline name should match ``target.kind.name`` (e.g. ``"cuda"``, | ||
| ``"hip"``, ``"c"``, ``"llvm"``). | ||
| """ | ||
| _PIPELINES[pipeline.name] = pipeline | ||
| return pipeline |
There was a problem hiding this comment.
Prevent silent pipeline overwrite on duplicate backend registration.
At Line 35, duplicate backend names overwrite prior registrations without any signal. That can silently switch the lowering path for a target and make behavior depend on import/registration order.
Suggested fix
def register_pipeline(pipeline: Pipeline) -> Pipeline:
@@
- _PIPELINES[pipeline.name] = pipeline
+ if pipeline.name in _PIPELINES:
+ raise ValueError(
+ f"Pipeline '{pipeline.name}' is already registered."
+ )
+ _PIPELINES[pipeline.name] = pipeline
return pipeline🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@tilelang/backend/pipeline.py` around lines 29 - 36, The register_pipeline
function currently overwrites existing entries in the _PIPELINES mapping when
pipeline.name already exists; change register_pipeline to first check if
pipeline.name is already present in _PIPELINES and, instead of silently
replacing it, raise a clear exception (e.g., ValueError) or log/error and refuse
registration so callers must explicitly unregister/replace before adding;
reference the register_pipeline function, the _PIPELINES mapping, and
pipeline.name when locating where to add the existence check and error handling.
There was a problem hiding this comment.
Actionable comments posted: 1
Caution
Some comments are outside the diff and can’t be posted inline due to platform limitations.
⚠️ Outside diff range comments (2)
tilelang/backend/cuda/pipeline.py (1)
20-29:⚠️ Potential issue | 🟡 Minor | ⚡ Quick winGuard against
Nonetarget inallow_warp_specialized(CUDA pipeline)
- In
tilelang/backend/cuda/pipeline.pylines 20-29,target: Target | Noneis forwarded tois_cuda_target(target)andhave_tma(target)without checking forNone.is_cuda_target(tilelang/jit/adapter/utils.py) directly readstarget.kind.name(noNonehandling), andhave_tma(tilelang/contrib/nvcc.py) similarly readstarget.kind.name(noNonehandling), sotarget=Nonewould raise.- Even if current call sites pass a non-
Nonetarget, the function’s signature implies it can be called without one—addif target is None: return False(and mirror intilelang/engine/phase.py) or tighten the signature to requireTarget.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tilelang/backend/cuda/pipeline.py` around lines 20 - 29, The function allow_warp_specialized currently forwards target to is_cuda_target and have_tma without guarding for None; add an early check "if target is None: return False" at the top of allow_warp_specialized so neither is_cuda_target nor have_tma are called with None, and mirror the same guard in the related entry (tilelang/engine/phase.py) or alternately change the signature to require a non-None Target; keep the existing pass_ctx handling and the config key "tl.disable_warp_specialized" unchanged.tilelang/engine/pass_pipeline.py (1)
24-27:⚠️ Potential issue | 🔴 Critical | ⚡ Quick winFix crash in
CommonPassPipelineBody:should_enable_aggressive_mergecalled with unsupportedtarget=
tilelang/engine/pass_pipeline.pydefinesshould_enable_aggressive_merge(pass_ctx=...)with notargetparameter, butCommonPassPipelineBodycalls it asshould_enable_aggressive_merge(pass_ctx=pass_ctx, target=target), which will raiseTypeErrorand abort lowering. (There’s a target-aware version of this helper intilelang/engine/phase.py.)Proposed minimal fix
- enable_aggressive_merge = should_enable_aggressive_merge(pass_ctx=pass_ctx, target=target) + enable_aggressive_merge = should_enable_aggressive_merge(pass_ctx=pass_ctx)🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tilelang/engine/pass_pipeline.py` around lines 24 - 27, CommonPassPipelineBody is calling should_enable_aggressive_merge(pass_ctx=pass_ctx, target=target) but the current helper signature def should_enable_aggressive_merge(pass_ctx: PassContext | None = None) lacks a target parameter and raises TypeError; update should_enable_aggressive_merge to accept an optional target parameter (e.g., target: str | None = None) and either forward to the target-aware implementation in tilelang.engine.phase or use target when selecting the config, ensuring the function still works when called without target and preserves existing behavior when pass_ctx is None by obtaining the PassContext via tilelang.transform.get_pass_context().
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@tilelang/backend/cuda/pipeline.py`:
- Around line 33-38: The docstring currently references the pass name
"LowerAndLegalize" but the pipeline actually calls LowerTileOp and relies on the
tl.has_tma attribute set by that pass; update the docstring to mention
LowerTileOp instead of LowerAndLegalize so it matches the implementation (the
check that reads tl.has_tma set by LowerTileOp). Ensure the text explicitly
names LowerTileOp and keeps the rest of the explanation intact.
---
Outside diff comments:
In `@tilelang/backend/cuda/pipeline.py`:
- Around line 20-29: The function allow_warp_specialized currently forwards
target to is_cuda_target and have_tma without guarding for None; add an early
check "if target is None: return False" at the top of allow_warp_specialized so
neither is_cuda_target nor have_tma are called with None, and mirror the same
guard in the related entry (tilelang/engine/phase.py) or alternately change the
signature to require a non-None Target; keep the existing pass_ctx handling and
the config key "tl.disable_warp_specialized" unchanged.
In `@tilelang/engine/pass_pipeline.py`:
- Around line 24-27: CommonPassPipelineBody is calling
should_enable_aggressive_merge(pass_ctx=pass_ctx, target=target) but the current
helper signature def should_enable_aggressive_merge(pass_ctx: PassContext | None
= None) lacks a target parameter and raises TypeError; update
should_enable_aggressive_merge to accept an optional target parameter (e.g.,
target: str | None = None) and either forward to the target-aware implementation
in tilelang.engine.phase or use target when selecting the config, ensuring the
function still works when called without target and preserves existing behavior
when pass_ctx is None by obtaining the PassContext via
tilelang.transform.get_pass_context().
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: 9cabcc67-b996-4519-9ac5-3e0652c54d54
📒 Files selected for processing (4)
tilelang/backend/common.pytilelang/backend/cuda/pipeline.pytilelang/backend/rocm/pipeline.pytilelang/engine/pass_pipeline.py
🚧 Files skipped from review as they are similar to previous changes (2)
- tilelang/backend/common.py
- tilelang/backend/rocm/pipeline.py
| """Check if any function in the module was lowered with TMA operations. | ||
|
|
||
| This reads the ``tl.has_tma`` attribute set by ``LowerTileOp`` during | ||
| ``LowerAndLegalize``, which is the source of truth for whether TMA | ||
| copies were actually generated. | ||
| """ |
There was a problem hiding this comment.
Update docstring to match actual pass name.
The docstring mentions LowerAndLegalize, but the actual pipeline calls LowerTileOp() at line 87. Update the docstring to reference the correct pass name.
📝 Proposed fix
- This reads the ``tl.has_tma`` attribute set by ``LowerTileOp`` during
- ``LowerAndLegalize``, which is the source of truth for whether TMA
- copies were actually generated.
+ This reads the ``tl.has_tma`` attribute set by ``LowerTileOp``,
+ which is the source of truth for whether TMA copies were actually
+ generated.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@tilelang/backend/cuda/pipeline.py` around lines 33 - 38, The docstring
currently references the pass name "LowerAndLegalize" but the pipeline actually
calls LowerTileOp and relies on the tl.has_tma attribute set by that pass;
update the docstring to mention LowerTileOp instead of LowerAndLegalize so it
matches the implementation (the check that reads tl.has_tma set by LowerTileOp).
Ensure the text explicitly names LowerTileOp and keeps the rest of the
explanation intact.
|
@regression-perf |
There was a problem hiding this comment.
Actionable comments posted: 1
🧹 Nitpick comments (1)
tilelang/backend/metal/pipeline.py (1)
23-28: 💤 Low valueConsider passing
pass_ctxconsistently to predicate functions.
pass_ctxis retrieved on line 21 but not passed toshould_force_let_inline()andshould_enable_race_check(), whileallow_vectorize(),should_enable_aggressive_merge(), andshould_disable_shared_memory_reuse()later in the function receive it explicitly. If these functions internally callget_pass_context()anyway, this is fine; otherwise, consider passingpass_ctxfor consistency.🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@tilelang/backend/metal/pipeline.py` around lines 23 - 28, pass_ctx is retrieved earlier but not forwarded to the predicate calls should_force_let_inline() and should_enable_race_check(); update those calls to accept and pass the already-obtained pass_ctx for consistency with allow_vectorize(pass_ctx), should_enable_aggressive_merge(pass_ctx), and should_disable_shared_memory_reuse(pass_ctx) so all predicate functions use the same context (i.e., replace should_force_let_inline() with should_force_let_inline(pass_ctx) and should_enable_race_check() with should_enable_race_check(pass_ctx)).
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Inline comments:
In `@tilelang/backend/pipeline_utils.py`:
- Around line 55-80: The get_layout_visual_formats function currently allows
"all" to appear inside comma-separated lists (e.g., "txt,all,png"); update the
validation so that after splitting into formats_list you explicitly reject "all"
when formats_list has more than one item by raising a ValueError (keep the
existing expansion behavior when formats_str == "all"); include a clear error
message referencing TL_LAYOUT_VISUALIZATION_FORMATS and that "all" must be used
alone, and keep the existing invalid_formats check for other invalid tokens (use
function name get_layout_visual_formats and variable formats_list to locate
where to add the extra check).
---
Nitpick comments:
In `@tilelang/backend/metal/pipeline.py`:
- Around line 23-28: pass_ctx is retrieved earlier but not forwarded to the
predicate calls should_force_let_inline() and should_enable_race_check(); update
those calls to accept and pass the already-obtained pass_ctx for consistency
with allow_vectorize(pass_ctx), should_enable_aggressive_merge(pass_ctx), and
should_disable_shared_memory_reuse(pass_ctx) so all predicate functions use the
same context (i.e., replace should_force_let_inline() with
should_force_let_inline(pass_ctx) and should_enable_race_check() with
should_enable_race_check(pass_ctx)).
🪄 Autofix (Beta)
Fix all unresolved CodeRabbit comments on this PR:
- Push a commit to this branch (recommended)
- Create a new PR with the fixes
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: ab2aa29b-da72-4347-9932-ed9dda584bb5
📒 Files selected for processing (13)
testing/python/transform/_transform_testing_utils.pytilelang/backend/__init__.pytilelang/backend/common.pytilelang/backend/cpu/__init__.pytilelang/backend/cpu/pipeline.pytilelang/backend/cuda/pipeline.pytilelang/backend/metal/__init__.pytilelang/backend/metal/pipeline.pytilelang/backend/pipeline_utils.pytilelang/backend/rocm/pipeline.pytilelang/engine/lower.pytilelang/engine/semantic_check.pytilelang/jit/adapter/utils.py
✅ Files skipped from review due to trivial changes (2)
- tilelang/backend/metal/init.py
- tilelang/backend/cpu/init.py
🚧 Files skipped from review as they are similar to previous changes (5)
- tilelang/backend/init.py
- tilelang/jit/adapter/utils.py
- tilelang/engine/lower.py
- testing/python/transform/_transform_testing_utils.py
- tilelang/backend/cuda/pipeline.py
| def get_layout_visual_formats(pass_ctx: PassContext | None = None) -> list[str]: | ||
| if pass_ctx is None: | ||
| pass_ctx = tilelang.transform.get_pass_context() | ||
| formats_value = pass_ctx.config.get(tilelang.PassConfigKey.TL_LAYOUT_VISUALIZATION_FORMATS, "") | ||
| if not formats_value: | ||
| return ["txt"] | ||
|
|
||
| formats_str = formats_value.strip().lower() | ||
| valid_formats = ["txt", "png", "pdf", "svg", "all"] | ||
|
|
||
| if formats_str == "all": | ||
| return ["txt", "png", "pdf", "svg"] | ||
|
|
||
| if "," in formats_str: | ||
| formats_list = [f.strip() for f in formats_str.split(",")] | ||
| else: | ||
| formats_list = [formats_str] | ||
|
|
||
| invalid_formats = [f for f in formats_list if f not in valid_formats] | ||
| if invalid_formats: | ||
| raise ValueError( | ||
| f"Invalid formats for TL_LAYOUT_VISUALIZATION_FORMATS: {invalid_formats}. " | ||
| f"Valid formats are: {valid_formats}. " | ||
| f"You can choose one of the valid formats or a comma-separated list of formats.(e.g., 'txt,png,pdf')" | ||
| ) | ||
| return formats_list |
There was a problem hiding this comment.
Validate that "all" is not mixed with other formats.
The validation logic allows "all" to appear in a comma-separated list (e.g., "txt,all,png"), but "all" should only be valid as a standalone value. When used alone, it expands to all formats (line 66), but when mixed with other formats in a list, it passes validation and is returned as-is, which the downstream LayoutVisual function likely cannot handle.
🛡️ Proposed fix to reject "all" in comma-separated lists
if formats_str == "all":
return ["txt", "png", "pdf", "svg"]
if "," in formats_str:
formats_list = [f.strip() for f in formats_str.split(",")]
else:
formats_list = [formats_str]
+ if "all" in formats_list:
+ raise ValueError(
+ "The format 'all' cannot be used in a comma-separated list. "
+ "Use 'all' alone to enable all formats, or specify individual formats (e.g., 'txt,png,pdf')."
+ )
+
invalid_formats = [f for f in formats_list if f not in valid_formats]📝 Committable suggestion
‼️ IMPORTANT
Carefully review the code before committing. Ensure that it accurately replaces the highlighted code, contains no missing lines, and has no issues with indentation. Thoroughly test & benchmark the code to ensure it meets the requirements.
| def get_layout_visual_formats(pass_ctx: PassContext | None = None) -> list[str]: | |
| if pass_ctx is None: | |
| pass_ctx = tilelang.transform.get_pass_context() | |
| formats_value = pass_ctx.config.get(tilelang.PassConfigKey.TL_LAYOUT_VISUALIZATION_FORMATS, "") | |
| if not formats_value: | |
| return ["txt"] | |
| formats_str = formats_value.strip().lower() | |
| valid_formats = ["txt", "png", "pdf", "svg", "all"] | |
| if formats_str == "all": | |
| return ["txt", "png", "pdf", "svg"] | |
| if "," in formats_str: | |
| formats_list = [f.strip() for f in formats_str.split(",")] | |
| else: | |
| formats_list = [formats_str] | |
| invalid_formats = [f for f in formats_list if f not in valid_formats] | |
| if invalid_formats: | |
| raise ValueError( | |
| f"Invalid formats for TL_LAYOUT_VISUALIZATION_FORMATS: {invalid_formats}. " | |
| f"Valid formats are: {valid_formats}. " | |
| f"You can choose one of the valid formats or a comma-separated list of formats.(e.g., 'txt,png,pdf')" | |
| ) | |
| return formats_list | |
| def get_layout_visual_formats(pass_ctx: PassContext | None = None) -> list[str]: | |
| if pass_ctx is None: | |
| pass_ctx = tilelang.transform.get_pass_context() | |
| formats_value = pass_ctx.config.get(tilelang.PassConfigKey.TL_LAYOUT_VISUALIZATION_FORMATS, "") | |
| if not formats_value: | |
| return ["txt"] | |
| formats_str = formats_value.strip().lower() | |
| valid_formats = ["txt", "png", "pdf", "svg", "all"] | |
| if formats_str == "all": | |
| return ["txt", "png", "pdf", "svg"] | |
| if "," in formats_str: | |
| formats_list = [f.strip() for f in formats_str.split(",")] | |
| else: | |
| formats_list = [formats_str] | |
| if "all" in formats_list: | |
| raise ValueError( | |
| "The format 'all' cannot be used in a comma-separated list. " | |
| "Use 'all' alone to enable all formats, or specify individual formats (e.g., 'txt,png,pdf')." | |
| ) | |
| invalid_formats = [f for f in formats_list if f not in valid_formats] | |
| if invalid_formats: | |
| raise ValueError( | |
| f"Invalid formats for TL_LAYOUT_VISUALIZATION_FORMATS: {invalid_formats}. " | |
| f"Valid formats are: {valid_formats}. " | |
| f"You can choose one of the valid formats or a comma-separated list of formats.(e.g., 'txt,png,pdf')" | |
| ) | |
| return formats_list |
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
In `@tilelang/backend/pipeline_utils.py` around lines 55 - 80, The
get_layout_visual_formats function currently allows "all" to appear inside
comma-separated lists (e.g., "txt,all,png"); update the validation so that after
splitting into formats_list you explicitly reject "all" when formats_list has
more than one item by raising a ValueError (keep the existing expansion behavior
when formats_str == "all"); include a clear error message referencing
TL_LAYOUT_VISUALIZATION_FORMATS and that "all" must be used alone, and keep the
existing invalid_formats check for other invalid tokens (use function name
get_layout_visual_formats and variable formats_list to locate where to add the
extra check).
There was a problem hiding this comment.
🧹 Nitpick comments (1)
testing/python/transform/test_tilelang_transform_inject_tcgen05_fence.py (1)
5-5: ⚡ Quick winKeep these tests on a bounded CUDA lowering helper.
Importing and invoking
CUDAPassPipelineBodyhere makes the assertions depend on every later CUDA backend pass, not just thetcgen05_ld/stlowering these tests are trying to validate. Since this PR explicitly allows backend pipelines to diverge, unrelated pass reordering can now break these tests even when the transform under test is still correct. Please route these through the narrower CUDA test helper that stops at the required lowering stage instead of the full backend pipeline.Also applies to: 121-122, 169-170
🤖 Prompt for AI Agents
Verify each finding against current code. Fix only still-valid issues, skip the rest with a brief reason, keep changes minimal, and validate. In `@testing/python/transform/test_tilelang_transform_inject_tcgen05_fence.py` at line 5, The tests import and invoke CUDAPassPipelineBody which runs the full CUDA backend pipeline and makes assertions fragile; replace the direct use/import of CUDAPassPipelineBody with the project’s bounded CUDA test helper that stops at the tcgen05_ld/st lowering stage (i.e., import and call the narrower CUDA lowering helper instead of CUDAPassPipelineBody), update the import and call sites accordingly, and make the same replacement for the other occurrences in this file so the tests only exercise the intended lowering pass.
🤖 Prompt for all review comments with AI agents
Verify each finding against current code. Fix only still-valid issues, skip the
rest with a brief reason, keep changes minimal, and validate.
Nitpick comments:
In `@testing/python/transform/test_tilelang_transform_inject_tcgen05_fence.py`:
- Line 5: The tests import and invoke CUDAPassPipelineBody which runs the full
CUDA backend pipeline and makes assertions fragile; replace the direct
use/import of CUDAPassPipelineBody with the project’s bounded CUDA test helper
that stops at the tcgen05_ld/st lowering stage (i.e., import and call the
narrower CUDA lowering helper instead of CUDAPassPipelineBody), update the
import and call sites accordingly, and make the same replacement for the other
occurrences in this file so the tests only exercise the intended lowering pass.
ℹ️ Review info
⚙️ Run configuration
Configuration used: defaults
Review profile: CHILL
Plan: Pro
Run ID: df2cb2bd-6f18-4c60-9d50-25f93bfe06d1
📒 Files selected for processing (2)
testing/python/transform/test_tilelang_transform_inject_tcgen05_fence.pytilelang/backend/cuda/pipeline.py
🚧 Files skipped from review as they are similar to previous changes (1)
- tilelang/backend/cuda/pipeline.py
|
@regression-perf |
Performance Regression Test ReportTriggered by: @SiriusNEO Results
Artifacts
|
|
@regression-perf |
Performance Regression Test ReportTriggered by: @SiriusNEO Results
Artifacts
|
Tracking Issue: #2115
This PR majorly refactors the pass pipeline part. In our design, each backend should have its own pipeline since there are too many hardware-specific passes for now. Currently the pipelines in different backends are similar, but they will be gradually different in the future.
Summary by CodeRabbit
Refactor
New Features
Tests